#include <chrono>
#include <stdio.h>
#include <stdint.h>
#include <fstream>
#include <cuda_runtime.h>
#include <cuda/pipeline>
#include <cooperative_groups.h>

#define DIV_UP(size, align_size) (((size) + (align_size) - 1) / (align_size))
#define ST(X)  uint64_t start##X = clock();
#define ET(X)  uint64_t end##X = clock();

#define CUDA_CHECK(err)                                                                            \
    do {                                                                                           \
        cudaError_t err_ = (err);                                                                  \
        if (err_ != cudaSuccess) {                                                                 \
            std::printf("CUDA error %d at %s:%d\n", err_, __FILE__, __LINE__);                     \
            throw std::runtime_error("CUDA error");                                                \
        }                                                                                          \
    } while (0)

static std::uint64_t now() {
  return std::chrono::duration_cast<std::chrono::milliseconds>
    (std::chrono::high_resolution_clock::now().time_since_epoch()).count();
}

static void dumpbin(void *cuda_mem, size_t sz, const std::string &path) {
  char * buf = new char[sz];
  cudaMemcpy(buf, cuda_mem, sz, cudaMemcpyDeviceToHost);
  {
    std::ofstream ofs(path, std::ios::out | std::ios::trunc | std::ios::binary);
    if (ofs.is_open()){
      ofs.write((const char *)buf, sz);
      // printf("# dump %s\n", path.c_str());
    }else{
      printf("open %s failed!", path.c_str());
    }
  }
}

static void loadbin(void *cuda_mem, size_t sz, const std::string &path) {
  char * buf = new char[sz];
  {
    std::ifstream ifs(path, std::ios::in | std::ios::binary);
    if (ifs.is_open()){
      ifs.read((char *)buf, sz);
      cudaMemcpy(cuda_mem, buf, sz, cudaMemcpyHostToDevice);
      // printf("# load %s\n", path.c_str());
    }else{
      printf("open %s failed!", path.c_str());
    }
  }
}

template<typename T>
__global__ void mem_set(T *src, int w, T value, T bias){
  const int bx = blockIdx.x;
  const int tx = threadIdx.x;
  const int bw = blockDim.x;
  
  int idx = bx*bw + tx;
  if(idx < w){
    if(idx % 2){
      src[idx] = value+bias;
    }else{
      src[idx] = value-bias;
    }
  }
}